其實已經有點不知道在發什麼了 但既然這邊是廢文集散地
就讓我再多發三天的廢文吧~
這幾天簡單來說，就是 SA（Simulated Annealing，模擬退火）這個 Heuristic 演算法本身 使用在LCN(least crossing number) 這個問體，把它做平行化，然後再把它加到 GPU 上面的方法。
那從演算法的設計，到後續放到 GPU 上面，我都把它整理在這裡了。今天是 Day 27，讓我們繼續看下去吧！
Day 26 結尾已經把帳單寫出來了：
CPU creates one proposal
GPU evaluates one proposal
CPU waits for result
CPU decides acceptance
repeat thousands of times
如果 crossing kernel 只要幾十 microseconds，每一步的 launch、synchronization 與 host-device 控制轉移，確實可能比計算本身還貴。
這件事我們做了。Day 26 的 state machine 整段留在 device：
propose
-> validate
-> evaluate
-> accept
-> commit
-> publish best
-> cooling
timed loop 裡的合約是：
H2D bytes = 0
D2H bytes = 0
host decisions = 0
對讀過 Day 20、Day 23、Day 25 的人，這不該是驚喜。中間結果能不寫出去就不要寫；沒變的 layout 不要複製；force iteration 能不回 Python 就不要回。今天只是把同一句話套到 SA 的控制流程。
真正需要盯的是 kernel 裡面這張 mapping：
one CUDA block
-> one shared current layout
chains_per_group = 32
-> 32 threads 同時提案
-> 選一個 winner
-> 做一次 accept / reject
參數名叫 chains。跑起來卻不是 32 條 Markov chain。
今天要回答的不是「要不要每步回 CPU」。今天要回答：
把 loop 留在 GPU 之後，
我們還在跑 Day 26 那條單 proposal SA 嗎？
Payload 很小。Proposal 只有 node id、old position、new position，objective 也只有 (K, n_K, Phi, C)。
小資料、高頻率，帳單不在 PCIe 頻寬，而在固定成本：
useful payload
a few integers
fixed cost
kernel launch
driver scheduling
stream ordering
synchronization
host wake-up
Resident kernel 把循序 loop 從 host 移進 device：
for (int iter = 0; iter < iterations_in_this_launch; ++iter) {
generate_proposals();
evaluate_proposals();
select_one_candidate();
accept_or_reject();
commit_current();
maybe_publish_best();
cool_temperature();
}
before
one host decision per step
frequent H2D / D2H
frequent synchronization
resident
zero host decisions inside timed loop
zero H2D / D2H inside timed loop
host only schedules work at batch boundaries
Day 25 已經對 force-directed layout 做過同一招：多輪 iteration 留在一個 kernel，positions 仍從 global memory 讀 tile，省的是 launch 和回 Python。今天不是新發明「loop 不要回 host」。
它也沒有把 step 0 到 step 9999 同時執行。單條 chain 仍然有順序：
current state at step t
-> proposal t
-> accept / reject
-> current state at step t + 1
Device residency 加速 sequential control；它沒有消除 Markov chain 的 dependency。跨 round 仍然一步接一步。變了的是：一步裡面看了幾個 proposals、怎麼決定要不要走。
「Resident」很容易讓人誤以為整張 graph 永遠放在 SRAM。Day 25 的小圖路徑確實把一份 positions 留在一個 block 的 shared memory 裡。今天這條路徑不是那個意思。
目前 flash_sa_resident_loop_kernel 的主要 state 放在 device global memory：
current_x / current_y
best_x / best_y
current_edge_crossings
best_edge_crossings
candidate_edge_crossings
current / best objective
RNG state
runtime and telemetry
Kernel 裡明確宣告的 shared memory 只有幾個 block-wide control flags：
__shared__ int selected_proposal;
__shared__ int accepted;
__shared__ int publish_best;
因此：
Day 25 小圖路徑
一份 positions 住在一個 block 的 shared memory
Day 27 resident SA
決策與 state 留在 GPU 上，資料仍在 DRAM
Global memory traffic 仍然存在。Coordinates、per-edge counts 與 candidate results 仍由 device memory 讀寫。拿掉的是 PCIe、driver 與 host synchronization，不是所有 DRAM access。
Chains Per Group 從第一次出現就該叫 Proposal Workers目前 kernel 的 mapping 是：
one CUDA block
-> one group state
one active thread in that block
-> one proposal worker
many iterations inside the block
-> one shared current-state trajectory
每一輪中：
if (tid < chains_per_group) {
proposal[tid] = make_proposal(current, rng[tid]);
evaluation[tid] = evaluate(current, proposal[tid]);
}
接著 thread 0 從所有 valid candidates 裡選出 lexicographic objective 最好的那一個：
selected = -1;
for (int p = 0; p < chains_per_group; ++p) {
if (!evaluation[p].valid) continue;
if (selected < 0 || evaluation[p].objective < evaluation[selected].objective) {
selected = p;
}
}
最後只有 selected candidate 進入一次 acceptance decision，整個 group 最多 commit 一個 move。
因此同一 group 內的 32 個 chains 並不是 32 條各有 current state 的 Markov chains：
Day 26 的一條 chain
sample one proposal
-> accept / reject
-> one transition
今天的一個 group round
32 workers 讀同一份 current
-> 產生 32 個 candidates
-> 選一個 winner
-> 做一次 transition
比較準確的名稱是 parallel proposal workers。參數仍叫 chains_per_group，是歷史名字，不是搜尋語意。
真正獨立的多條 replicas 需要各自保存 current layout、objective、temperature 與 RNG stream。那是 Day 28 的主題。
這是今天最容易看走眼的地方。
Day 26 的單 proposal SA 每一步是：
sample one
-> legal?
-> Metropolis against current
-> maybe commit
Resident group round 是：
sample W
-> drop invalid
-> choose the lexicographic best among valid
-> Metropolis that winner against current
-> maybe commit once
兩個搜尋的 transition distribution 不一樣。
第一，每輪看到「好 proposal」的機會變高。W 個獨立樣本裡的最佳者，期望上優於只抽一個。高溫時這像同時探了 W 個鄰居；低溫時這像每步都先做一次 greedy filter，再決定要不要走。
第二，沒被選上的 valid proposal 從來沒有機會跟 current 做 Metropolis。Day 26 裡它們各自會走一次 accept / reject；今天它們在 winner scan 就被丟了。
第三，比較速度與解品質時，不能把 W 個 candidate evaluations 直接當成 W 個 sequential SA transitions：
candidate evaluations != SA transitions
SA transitions != best updates
best updates != K improvements
W = 32、跑 10,000 rounds，是 320,000 次 evaluation，但最多 10,000 次 chain transition。把前者除以後者，會得到一個看起來很勤快的搜尋，實際上 current 只走了 10,000 步，而且每步的提案分布已經被 winner filter 改過。
Thread 0 的 winner scan 本身也是 serial O(W)。當 proposal evaluation 很昂貴時，這段成本很小；若後續 spatial evaluator 已經把評分大幅壓低，它可能成為下一個需要改成 block reduction 的區段。最佳化後必須重新 profile，不能假設原本的小成本會永遠很小。
Day 26 的 current、best，搬到 GPU 後還需要一份 candidate work area。因為現在同一輪有 W 個 workers 同時寫自己的 scratch，不能再靠 Python 那份 list(current_positions) 的複製品。
| State | 內容 | 何時修改 |
|---|---|---|
| Current | coordinates、edge counts、objective | 唯一的 winner 被接受後 |
| Best | coordinates、edge counts、objective | current 嚴格優於 best 時 |
| Candidate | 每個 worker 的 proposal、objective、edge counts | 每輪覆寫 |
| Control | temperature、iteration、plateau age | 每輪結束，整個 group 共用一份 |
| RNG | 每個 worker 的亂數狀態 | 產生 proposal 與 acceptance draw |
| Telemetry | attempts、valid、accepted、branches、best updates | 對應事件發生時 |
原專案的 workspace 直接預先配置：
current_x, current_y
best_x, best_y
scratch_x, scratch_y
current_edge_crossings
best_edge_crossings
scratch_edge_crossings
current_objective
best_objective
scratch_objective
另外還有每個 candidate worker 的 proposal、evaluation result 與 edge-count buffer。
這些 allocations 在 timed loop 之前完成。Resident tests 明確要求：
timed_cuda_malloc_count = 0
timed_cuda_free_count = 0
因為 cudaMalloc 不只是取得記憶體，也可能帶來昂貴的 runtime coordination。把 allocation 搬出 hot loop，通常比替一個很小的 device function 手動展開幾行更有價值。
注意 temperature 是 group 共用的。32 個 workers 沒有 32 個溫度。這又是「不是 32 條 chain」的一個具體後果：它們甚至不能各自冷卻。
不必把 Day 26 的 SA 再講一次。這裡只看同一輪裡，哪些動作是 W 份平行、哪些動作仍然只有一次。
每個 worker 有自己的 RNG：
rng[group, worker]
固定模式會選 node，再依目前溫度決定 local radius，寫一份 proposal。Proposal 保存 from_x、from_y。若 current 已經不在那個 old position，這份 proposal 對錯版本的 state，直接標成 invalid。
Day 26 會複製整張 candidate layout。GPU 改用 virtual override，這是 Day 23 同一招：
candidate_x(node) =
node == moved_node ? proposal.to_x : current_x[node];
full candidate coordinate copy
O(N) writes per proposal
virtual override
O(1) proposal state
每個 proposal 先過 V1～V4。非法的不進 winner scan：溫度只能接受較差的合法解，不能接受 geometry violation。
活下來的 worker 才算 candidate edge counts 與 (K, n_K, Phi, C)。Full-square 從頭算全部 pair；incremental 從 current_edge_crossings 出發，只處理 moved node 的 incident edges。若 sum(count[e]) 是奇數，incremental closing state 已經壞了，evaluator 要求 fallback，不能帶著損壞的 counts 繼續跑。
到這裡為止，平行度來自多個 proposals，不是讓整個 block 合作評估一個 proposal。
所有 workers 完成 evaluation 後：
__syncthreads();
Thread 0 選 winner，再套用 Day 26 的 layered acceptance。接受後，thread 0 先更新 moved node 與 objective，整個 block 再合作複製 selected candidate 的 per-edge counts：
current_x[node] = proposal.to_x;
current_y[node] = proposal.to_y;
current_objective = selected.objective;
for (int edge = tid; edge < E; edge += blockDim.x) {
current_edge_crossings[edge] = selected_edge_crossings[edge];
}
Rejected 完全不寫 current。這就是 transactional scratch：先在 candidate buffer 形成完整、一致的候選狀態，接受後才發布成 current。
接受也不等於更新 best：
publish_best = accepted && current_objective < best_objective;
只有這時才複製 current → best。最後 thread 0 更新 iteration、plateau age 與 temperature。整個 group 共用這一份控制狀態。
所以一輪結束時，current 最多走一步。W 只改變「這一步看過哪些鄰居」，不改變「這一步能走幾次」。
一輪裡的重要同步點是：
all proposals evaluated
|
barrier
v
thread 0 selects and decides acceptance
|
barrier
v
all threads know accepted and selected index
|
copy selected counts to current
|
barrier
v
thread 0 decides publish_best
|
barrier
v
copy current state to best if needed
|
barrier
v
update temperature and start next iteration
如果缺少 selected-count copy 後的 barrier，下一輪可能看到：
new coordinates
new objective
partly old edge counts
這份 state 表面上每個欄位都存在，彼此卻描述不同的 layout。Incremental evaluator 會從錯誤 cache 繼續更新，問題可能過很多輪才被 sum(count[e]) 奇偶性或 exact validator 發現。
Winner-then-Metropolis 已經改了搜尋。Barrier 則保證：改過之後的那一份 current，仍然是同一世代的 coordinates、counts 與 objective。Synchronization 在這裡不是單純的效能損失，它界定 commit transaction 的可見性。
Reject 確實避免了兩份發布成本：
no current edge-count copy
no best coordinate / count copy
但每個 worker 在被丢掉之前已經做了：
geometry validation
crossing evaluation
candidate objective construction
candidate edge-count writes
若每個 group 有 W 個 workers，每個 worker 保存 E 個 candidate counts，單輪 candidate scratch 至少是：
candidate_count_bytes = W * E * 4
例如：
W = 256
E = 100,000
bytes = 256 * 100,000 * 4
= 102,400,000 bytes
= about 97.66 MiB per group
這還沒算 coordinates、current/best counts、objectives、spatial index 與其他 work buffers。
所以「同時產生更多 proposals」會增加平行度，也會線性增加 candidate state。它同時在兩件事上收費：每輪多做的 evaluation，以及那份永遠用不到的 loser scratch。當 evaluator 已經很快時，global-memory footprint 和寫入 candidate counts 可能成為新的瓶頸。
進一步的方向是只保存 incident-edge deltas，接受時再套用到 current counts；代價是 commit 與 rollback contract 會更複雜，而且必須保證沒有漏掉 static edge 被 incident edge 新增或移除的 crossing contribution。
Proposals Per Launch 不是 Concurrent Proposal 數這裡有兩個很像、但意義不同的參數：
chains_per_group
同一輪同時評估多少 proposals
也就是 proposal worker 數
它改變搜尋語意
proposals_per_launch
一次 kernel launch 在同一 current trajectory 上執行多少輪
它只改變控制粒度
它不改變每一輪仍只 commit 一次
假設：
iterations = 10,000
proposals_per_launch = 4
Resident loop kernel launch 數是：
ceil(10,000 / 4) = 2,500 launches
如果每一步至少需要一個 host-controlled evaluation launch，單看 resident loop launch 可得到：
10,000 / 2,500 = 4x fewer loop launches
這仍不是 4 倍 runtime speedup，更不是 4 倍搜尋進度。每個 batch 後還可能有 group-sync kernel，setup 與 final hash 也有額外 launches，而真正的 evaluator 工作量沒有因 batching 消失。每一輪仍然最多走一步。
把 proposals_per_launch 拉大可以攤平 launch overhead，但會降低控制粒度：
larger batch
fewer launches
longer interval before host regains control
slower response to time limit or external policy changes
group exchange cannot occur inside the batch boundary
若 sync_interval 比 batch 更早到，host driver 會縮短這次 launch，確保 group synchronization 發生在指定 iteration。Day 28 的 replica exchange 就卡在這個邊界上：exchange 不能發生在 batch 內部，因為 host 當時看不見任何 replica state。
Day 27 加了一個不需要 GPU 的模型：
python case_4/examples/sa_residency_model.py
預設參數是：
N = 150
E = 300
groups = 1
workers per group = 32
iterations = 10,000
proposals per launch = 4
輸出：
candidate evaluations: 320,000
minimum host evaluation launches: 10,000
resident loop launches: 2,500
resident launch reduction: 4.00x
coordinate states / group: 3.52 KiB
edge-count states / group: 3.52 KiB
candidate edge counts / group: 37.50 KiB
modeled lower bound / group: 44.53 KiB
請先看第一行和「10,000 iterations」的關係：
320,000 evaluations
/ 10,000 rounds
= 32 workers 同時看同一份 current
最多 10,000 次 transition
因為每輪只 commit 一次
Memory lower bound 使用：
three coordinate states
current + best + scratch
3 * 2 * N * 4 bytes
three edge-count states
current + best + scratch
3 * E * 4 bytes
per-worker candidate counts
W * E * 4 bytes
它沒有包含 allocator metadata、objectives、RNG、telemetry、spatial grid、validator 與其他 candidate buffers，因此只能叫做 lower bound。
同樣地，launch reduction 是控制流算術，不是 GPU benchmark，也不是搜尋進度。
在 host-driven path 中，profiler 可能顯示大量：
short kernels
cudaMemcpy
cudaStreamSynchronize
CPU gaps between launches
Resident 化之後，這些固定成本下降，真正的計算才會成為主角。也是在這個時候，才看得出今天這條 chain 的平行度有多窄。
目前 resident evaluator 的工作 mapping 是：
one worker thread
-> validate one proposal
-> loop nodes and edges
-> build one candidate objective
如果 chains_per_group = 32，256-thread block 中只有前 32 個 threads 進行 proposal evaluation，剛好是一個 warp。其餘 threads 要等到 cooperative copy 或 barrier 階段才可能工作。
GPU 並沒有在「加速同一條 Day 26 chain 的單步計算」。它把多餘的 threads 拿去同時產更多 candidates，然後再丟掉大多數。
Full-square evaluator 的單輪幾何工作量接近：
W * E * E crossing predicates
從 Day 23 的演算法目標來看，incremental crossing predicates 應該接近：
W * degree(moved_node) * E
但目前 resident function 仍用兩層 for edge / for other 掃過 E * E 個 pair slots，再用 incident 判斷跳過不受影響的 pair：
pair-loop visits
W * E * E
expensive intersection predicates
approximately W * degree(moved_node) * E
所以它確實減少 orientation 與 segment-intersection 計算，卻沒有消除所有平方級 loop、branch 與 edge-endpoint reads。真正的 spatial candidate list 應該連不可能相交的 partner indices 都不要走訪。
當 launch 與 host round-trip 不再遮住它，Day 23 的增量化才會直接決定吞吐量。瓶頸會移動：
remove boundary overhead
-> expose evaluator cost
remove unnecessary pair work
-> expose memory traffic
compress candidate state
-> expose selection and synchronization
使用 persistent kernel 不會讓瓶頸永遠消失。它只是讓你終於看到：這顆 kernel 大部分時間在為不會被 commit 的 proposals 工作。
PreProposalEvaluatorMode 列出：
FULL_SQUARE_GPU
INCREMENTAL_GPU
INCREMENTAL_C1_DENSE
INCREMENTAL_C2_CLOSING_STATE
INCREMENTAL_C3_SPATIAL
但目前快照中的 flash_resident_evaluate(...) 只把模式分成兩類：
full_square = mode == FULL_SQUARE_GPU;
if full_square:
recompute all pairs
else:
use the same incident-edge incremental branch
也就是這個 resident function 本身沒有針對 C1、C2、C3 選擇不同的 specialized implementation。其他 replay 或 qualification paths 有各自的實作，不代表 resident hot loop 自動使用相同的 spatial candidate list。
因此 benchmark 不能只根據 CLI 顯示 INCREMENTAL_C3_SPATIAL，就把效能歸因給 C3。應該用 profiler、kernel symbol、工作量 counters，或刻意設計能區分 C1/C3 掃描量的測試，確認實際執行路徑。
這個例子非常典型：
configuration name
does not prove
executed optimization
它和 chains_per_group 是同一類陷阱：標籤看起來像你想要的演算法，跑起來是另一個。
Rejected 也要小心解讀每輪 W 個 workers 產生 candidates，最多一個會被 commit。Kernel 的 counter 使用：
rejected_count += W - (accepted ? 1 : 0)
因此 rejected 包含：
invalid proposals
valid but not selected proposals
selected proposal rejected by Metropolis
若 W = 256，即使每輪 winner 都被接受：
accepted = 1
rejected = 255
表面 acceptance rate 只有：
1 / 256
這不能解讀成溫度太低或 Metropolis 幾乎不接受。大部分「rejected」其實連 acceptance branch 都沒有進入。它們是 winner-then-Metropolis 的直接後果，不是 Day 26 那條 chain 的拒答率。
所以還要同時看：
candidate attempts
geometry-valid candidates
selected candidates
Metropolis accept / reject branches
committed moves
best updates
沒有分母定義的 acceptance rate，很容易比較錯誤。尤其是拿它去跟 Day 26 的單 proposal telemetry 對打：兩邊的 rejected 根本不是同一個事件。
這裡必須講清楚測了什麼、沒測什麼。
Residency tests 證明的是 state、residency 與 deterministic contract。它們不證明「這還是 Day 26 那條 chain」——因為它本來就不是。Winner-then-Metropolis 是刻意改過的搜尋。測試要守住的是：改過之後，current / best 仍然合法、可重播、而且 timed loop 真的沒有跑回 host。
第一，group 與 worker width：
groups = 1, 2, 4
workers per group = 1, 2, 32, 256
每組的 best objective 必須不差於 current objective：
best <= current
而且 published state 必須 exact_valid = 1。
第二，固定 seed 的 replay 必須 deterministic：
same config + same seed
-> same current state hash
-> same best state hash
-> same objectives
測試涵蓋 full-square 與所有 incremental mode labels。
第三，temperature 必須在 device 更新。10 iterations 後：
actual temperature
approximately equals
0.1 * 0.999 ** 10
第四，timed loop 的 residency counters 必須是：
cudaMalloc count = 0
cudaFree count = 0
H2D bytes = 0
D2H bytes = 0
host decisions = 0
第五，非法 worker width 在 launch 前拒絕：
chains_per_group > 256
-> configuration error
這些測試不自動證明 C3 比 C1 快，不提供端到端 speedup，也不證明 W = 32 的搜尋質量等於 32 條 Day 26 chain。W = 1 那一組才最接近昨天的單 proposal 語意；其餘寬度是另一個演算法。
目前快照保留 residency runner 與測試，沒有保留能直接引用的同硬體、同 commit、host-driven 對 resident 的正式 timing table。因此這篇不填一個猜測的 speedup，也不用「少了每步回 CPU」去暗示已經很快。
需要分三層測量：
control-plane evidence
H2D / D2H bytes
host decision count
kernel launch count
synchronization count
kernel evidence
duration per resident launch
crossing predicates per second
global load/store throughput
branch divergence
eligible warps and achieved occupancy
search evidence
proposals per second
candidate evaluations per second
committed moves per second
time to target K
final validated K at fixed wall-clock budget
其中最重要的單位區別仍然是：
candidate evaluations != SA transitions
Resident path 可能讓 candidate evaluations / second 大幅上升，卻因 winner-selection policy，沒有同比改善 time-to-target-K。那不是量錯，那是今天改過的 chain 在說話。
Day 25 的 Layout 透過 shared-memory tiling 減少重複讀取，也把多輪 force iteration 留在 device。Day 27 看起來像在處理更上層的系統成本：
keep control and state on device
-> reduce launch boundaries
-> remove per-step H2D / D2H
-> remove host decisions
對這個系列，那句話已經不夠當主軸。更值錢的是：把控制留在 GPU 之後，平行提案會順便改掉搜尋。
具體技術包括：
也有幾個反直覺結果：
resident does not mean shared-memory resident
32 chains_per_group is not 32 Markov chains
sample W then Metropolis one
is not W sequential SA steps
more workers means more candidate memory
and more evaluations that will never commit
larger proposals_per_launch reduces launches
but also reduces control responsiveness
configuration name
does not prove specialized code executed
rejected telemetry is not Metropolis rejection rate
Day 28 會把一個 shared-current group 擴展成真正的多 replica 搜尋。每個 replica 有自己的 layout 與 temperature，GPU 可以讓它們同時探索；難題則變成 replicas 何時交換、交換的是 state 還是 temperature label，以及 acceptance probability 要如何保持正確。
那才是「同時跑很多條 SA」。今天這 32 個 workers，還只是同一條 current 上的平行提案。
iThome鐵人賽
看影片追技術